#pragma OPENCL EXTENSION cl_arm_printf : enable
__kernel void matmul_fp32_lbrbobntnp(__global float *out_buf, __global float4 *lhs_buf,
                                   __global float4 *rhs_buf, int4 hwkt) {
    int height = hwkt.x, width = hwkt.y, kslice = hwkt.z/4;
    int h = get_global_id(0);
    int w = get_global_id(1);    

    // 778 ms
    // float acc = 0;
    // for(int i=0; i<kslice; i++) {
    //     acc += lhs_buf[h*kslice + i + 0].x * rhs_buf[w*kslice + i + 0].x;
    //     acc += lhs_buf[h*kslice + i + 0].y * rhs_buf[w*kslice + i + 0].y;
    //     acc += lhs_buf[h*kslice + i + 0].z * rhs_buf[w*kslice + i + 0].z;
    //     acc += lhs_buf[h*kslice + i + 0].w * rhs_buf[w*kslice + i + 0].w;
    // }
    // out_buf[h*width + w] = acc;

    // 684 ms
    // float acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc += lhs_buf[h*kslice + i + 0].x * rhs_buf[w*kslice + i + 0].x;
    //     acc += lhs_buf[h*kslice + i + 0].y * rhs_buf[w*kslice + i + 0].y;
    //     acc += lhs_buf[h*kslice + i + 0].z * rhs_buf[w*kslice + i + 0].z;
    //     acc += lhs_buf[h*kslice + i + 0].w * rhs_buf[w*kslice + i + 0].w;

    //     acc += lhs_buf[h*kslice + i + 1].x * rhs_buf[w*kslice + i + 1].x;
    //     acc += lhs_buf[h*kslice + i + 1].y * rhs_buf[w*kslice + i + 1].y;
    //     acc += lhs_buf[h*kslice + i + 1].z * rhs_buf[w*kslice + i + 1].z;
    //     acc += lhs_buf[h*kslice + i + 1].w * rhs_buf[w*kslice + i + 1].w;

    //     acc += lhs_buf[h*kslice + i + 2].x * rhs_buf[w*kslice + i + 2].x;
    //     acc += lhs_buf[h*kslice + i + 2].y * rhs_buf[w*kslice + i + 2].y;
    //     acc += lhs_buf[h*kslice + i + 2].z * rhs_buf[w*kslice + i + 2].z;
    //     acc += lhs_buf[h*kslice + i + 2].w * rhs_buf[w*kslice + i + 2].w;

    //     acc += lhs_buf[h*kslice + i + 3].x * rhs_buf[w*kslice + i + 3].x;
    //     acc += lhs_buf[h*kslice + i + 3].y * rhs_buf[w*kslice + i + 3].y;
    //     acc += lhs_buf[h*kslice + i + 3].z * rhs_buf[w*kslice + i + 3].z;
    //     acc += lhs_buf[h*kslice + i + 3].w * rhs_buf[w*kslice + i + 3].w;
    // }
    // out_buf[h*width + w] = acc;

    // 684 ms
    // float acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc = mad(lhs_buf[h*kslice + i + 0].x, rhs_buf[w*kslice + i + 0].x, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 0].y, rhs_buf[w*kslice + i + 0].y, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 0].z, rhs_buf[w*kslice + i + 0].z, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 0].w, rhs_buf[w*kslice + i + 0].w, acc);

    //     acc = mad(lhs_buf[h*kslice + i + 1].x, rhs_buf[w*kslice + i + 1].x, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 1].y, rhs_buf[w*kslice + i + 1].y, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 1].z, rhs_buf[w*kslice + i + 1].z, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 1].w, rhs_buf[w*kslice + i + 1].w, acc);

    //     acc = mad(lhs_buf[h*kslice + i + 2].x, rhs_buf[w*kslice + i + 2].x, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 2].y, rhs_buf[w*kslice + i + 2].y, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 2].z, rhs_buf[w*kslice + i + 2].z, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 2].w, rhs_buf[w*kslice + i + 2].w, acc);

    //     acc = mad(lhs_buf[h*kslice + i + 3].x, rhs_buf[w*kslice + i + 3].x, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 3].y, rhs_buf[w*kslice + i + 3].y, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 3].z, rhs_buf[w*kslice + i + 3].z, acc);
    //     acc = mad(lhs_buf[h*kslice + i + 3].w, rhs_buf[w*kslice + i + 3].w, acc);
    // }
    // out_buf[h*width + w] = acc;

    // 684 ms
    float4 acc = 0;
    for(int i=0; i<kslice; i+=4) {
        acc = mad(lhs_buf[h*kslice + i + 0], rhs_buf[w*kslice + i + 0], acc);
        acc = mad(lhs_buf[h*kslice + i + 1], rhs_buf[w*kslice + i + 1], acc);
        acc = mad(lhs_buf[h*kslice + i + 2], rhs_buf[w*kslice + i + 2], acc);
        acc = mad(lhs_buf[h*kslice + i + 3], rhs_buf[w*kslice + i + 3], acc);
    }
    out_buf[h*width + w] = acc.x + acc.y + acc.z + acc.w;

    // 684 ms
    // float acc = 0;
    // for(int i=0; i<kslice; i+=4) {
    //     acc += dot(lhs_buf[h*kslice + i + 0], rhs_buf[w*kslice + i + 0]);
    //     acc += dot(lhs_buf[h*kslice + i + 1], rhs_buf[w*kslice + i + 1]);
    //     acc += dot(lhs_buf[h*kslice + i + 2], rhs_buf[w*kslice + i + 2]);
    //     acc += dot(lhs_buf[h*kslice + i + 3], rhs_buf[w*kslice + i + 3]);
    // }
    // out_buf[h*width + w] = acc;
}